#ifndef VORONOI_FUNCTION
#define VORONOI_FUNCTION
#include "stdafx.cuh"



extern "C" bool cuPrintInit();
extern "C" void initCUDA();
extern "C" void voronoiClean();
__global__ void vecAdd(float* A , float * B , float *C);
__global__ void voronoi_kernel(float2* d_sites ,int* d_sitesLength,Event* d_eventHeap,float2* d_vertexPool,int* d_vertexPoolLength,float2* d_center,ArcHeapEle* d_arcHeap,EllipseArc* d_arcPool , int* d_arcPoolLength);
__device__ void removeOutofdateEvents(Event* d_eventHeap,int& d_eventHeapSize);

extern "C" bool cuPrintInit()  
{  
	cudaError_t err = cudaPrintfInit();  
	if(0 != strcmp("no error", cudaGetErrorString(err))) {
		puts("device is not ready!\n");
		exit(EXIT_FAILURE);
	}
	return true;
}  

extern "C" void voronoiClean()
{


}


//void 
extern "C" void initCUDA(){

	//init device		
	cudaDeviceProp deviceProp;
	cudaGetDeviceProperties(&deviceProp,0);
	if(!deviceProp.canMapHostMemory){
		fprintf(stderr,"Device %d cannot map host memory!\n",0);
		exit(EXIT_FAILURE);
	}
	cudaSetDeviceFlags(cudaDeviceMapHost);
	cuPrintInit();



	std::map<std::string,float2>sitesInt;
	std::ostringstream outStr;
	int PrecisionNum = 100;
	int DenseControl = 4;
	int pointNum = 8;
	radiusVaring = 0.0f;


	for( int i = 0; i < pointNum; ++i){
		outStr.str("");
		float2 pt;
		pt.x = rand() % (PrecisionNum + 1);
		pt.x -= fmod((float)pt.x , (float)DenseControl);
		pt.y = rand() % (PrecisionNum + 1);
		pt.y -= fmod((float)pt.y , (float)DenseControl);
		outStr << pt.x << " " << pt.y;
		//std::cout << outStr.str();
		sitesInt[outStr.str()] = pt;
	}


	int sitesMemsize = (((sitesInt.size()*sizeof(float2)+PAGE_SIZE-1)/PAGE_SIZE)*PAGE_SIZE);
	int sitesLengthMemsize = (((sizeof(int)+PAGE_SIZE-1)/PAGE_SIZE)*PAGE_SIZE);
	printf("sitesMemsize siteslengthmemsize %d %d\n",sitesMemsize,sitesLengthMemsize);

	h_sites=(float2*)VirtualAlloc(NULL,sitesMemsize,MEM_COMMIT,PAGE_READWRITE);
	h_sitesLength=(int*)VirtualAlloc(NULL,sitesLengthMemsize,MEM_COMMIT,PAGE_READWRITE);
	h_vertexPool = (float2*)malloc(VERTEX_POOL_SIZE*sizeof(float2));
	h_vertexPoolLength = (int*)malloc(sizeof(int));
	h_center = (float2*)malloc(sizeof(float2));
	h_arcsPool = (EllipseArc*)malloc(sizeof(EllipseArc) * ARC_POOL_SIZE);
	h_arcsPoolLength = (int*)malloc(sizeof(int));

	cutilSafeCall( cudaMalloc((void**)&d_eventHeap,sizeof(Event) * EVENT_HEAP_SIZE) );
	cutilSafeCall( cudaMalloc((void**)&d_vertexPool,sizeof(float2) * VERTEX_POOL_SIZE) );
	cutilSafeCall( cudaMalloc((void**)&d_vertexPoolLength,sizeof(int)));
	cutilSafeCall( cudaMalloc((void**)&d_center,sizeof(float2)) );
	cutilSafeCall( cudaMalloc((void**)&d_arcHeap,sizeof(ArcHeapEle)*ARC_HEAP_SIZE));
	cutilSafeCall( cudaMalloc((void**)&d_arcsPool,sizeof(EllipseArc)*ARC_POOL_SIZE));
	cutilSafeCall( cudaMalloc((void**)&d_arcsPoolLength,sizeof(int)));

	cutilSafeCall( cudaHostRegister(h_sites,sitesMemsize,cudaHostRegisterMapped) );
	cutilSafeCall( cudaHostRegister(h_sitesLength,sitesMemsize,cudaHostRegisterMapped) );


	h_sitesLength[0] = 0;
	for(std::map<std::string,float2>::iterator it = sitesInt.begin();it!=sitesInt.end();++it){
		float2 pt;
		pt = make_float2(it->second.x/float(PrecisionNum) - .5f,it->second.y / float(PrecisionNum) -0.5 );
		if( !(fabs(pt.x-h_center[0].x) +fabs(pt.x-h_center[0].y) < WXN_EPSILON)){
			//sites.push_back(pt);
			h_sites[h_sitesLength[0]++]=pt;
			printf( "initial Point pt.x %f pt.y %f\n",pt.x,pt.y);
			//std::cout << pt.x << " " << pt.y << "\n";
		}
	}
	printf("before %f\n",h_sites[0].x);
	printf("length %d\n" , h_sitesLength[0]);
	cudaHostGetDevicePointer((void**)&(d_sites),(void*)(h_sites),0);
	cudaHostGetDevicePointer((void**)&(d_sitesLength),(void*)h_sitesLength,0);

	int t1 = clock();
	voronoi_kernel<<<1,1>>>(d_sites,d_sitesLength,d_eventHeap,d_vertexPool,d_vertexPoolLength,d_center,d_arcHeap,d_arcsPool,d_arcsPoolLength);
	cutilCheckMsg("kernel launch failure");
#ifdef _DEBUG
	cutilSafeCall( cudaDeviceSynchronize() );
#endif
	cutilSafeCall( cudaMemcpy(h_vertexPool,d_vertexPool,sizeof(float2) * VERTEX_POOL_SIZE ,cudaMemcpyDeviceToHost ) );
	cutilSafeCall( cudaMemcpy(h_vertexPoolLength,d_vertexPoolLength,sizeof(int),cudaMemcpyDeviceToHost ) );
	cutilSafeCall( cudaMemcpy(h_center,d_center,sizeof(float2),cudaMemcpyDeviceToHost ) );
	cutilSafeCall( cudaMemcpy(h_arcsPool,d_arcsPool,sizeof(EllipseArc)*ARC_POOL_SIZE,cudaMemcpyDeviceToHost));
	cutilSafeCall( cudaMemcpy(h_arcsPoolLength,d_arcsPoolLength,sizeof(int),cudaMemcpyDeviceToHost));


	printf("time %f\n" , (float)(clock() - t1) / (float)(CLOCKS_PER_SEC) );
	cudaPrintfDisplay(stdout, true);
	printf("after %f\n" ,h_sites[0].x );
	printf("length %d\n", h_sitesLength[0]);

	cudaPrintfEnd();  

}


__global__ void vecAdd(float* A , float * B , float *C){
	int i = threadIdx.x;
	C[i] = A[i] + B[i];
	//printf("Kernel : A[%d]=%f,B[%d]=%f,C[%d]=%f\n",i,A[i],i,B[i],i,C[i]);
}



__global__ void voronoi_kernel(float2* d_sites ,int* d_sitesLength,Event* d_eventHeap,float2* d_vertexPool,int* d_vertexPoolLength,float2* d_center,ArcHeapEle* d_arcHeap,EllipseArc* d_arcsPool , int* d_arcsPoolLength){
	//d_sitesLength[0]++;
	//d_sites[0].x += 1;
	cuPrintf("kernel's block :%d\n", blockDim.x);  
	cuPrintf("kernel's thread :%d\n", blockIdx.x);  
	cuPrintf("kernel's thread :%d\n", blockDim.x*blockIdx.x+threadIdx.x);  
	




	int threadTotalId = blockDim.x * blockIdx.x + threadIdx.x;
	cuPrintf( "kernel %d\n" , threadTotalId );
	
	if( threadTotalId == 0 ){
		int d_eventHeapSize;
		int d_arcHeapSize;
		float d_distance;
		float d_distanceGap;
		EllipseArc* headArc;
		
		
		d_distance = 0;
		heapInitial<Event>(d_eventHeap,d_eventHeapSize);
		heapInitial<ArcHeapEle>(d_arcHeap,d_arcHeapSize);
		d_vertexPoolLength[0] = 0;
		d_center[0].x = d_center[0].y = 0;
		d_arcsPoolLength[0] = 0;
		headArc = NULL;
		

		Event evt;
		for(int i = 0; i < d_sitesLength[0];++i){
			evt.type = NewEllipse;
			evt.site = d_sites[i];
			evt.distance = sqrtf((d_center[0].x-d_sites[i].x)*(d_center[0].x-d_sites[i].x)+(d_center[0].y-d_sites[i].y)*(d_center[0].y-d_sites[i].y));		
			heapInsert<Event>(evt,d_eventHeap,d_eventHeapSize);
		}
		d_distanceGap = 0.01;
		if( !heapEmpty<Event>(d_eventHeapSize)&&heapTop<Event>(d_eventHeap)->distance > d_distanceGap ){
			d_distanceGap = heapTop<Event>(d_eventHeap)->distance;
		}
		cuPrintf("eventtop %f\n" , heapTop<Event>(d_eventHeap)->distance);
		
		//main step
		for(int i = 0; i < 10;++i){
			removeOutofdateEvents(d_eventHeap,d_eventHeapSize);
			if( heapEmpty<Event>(d_eventHeapSize) ) break;
			d_distance = heapTop<Event>(d_eventHeap)->distance;
			if( headArc == NULL){
				float dOfFirstLevel = heapTop<Event>(d_eventHeap)->distance;
				heapInitial<ArcHeapEle>(d_arcHeap,d_arcHeapSize);
				Event lastEvent;
				lastEvent = *(heapTop<Event>(d_eventHeap));

				int cnt = 0;
				while(!heapEmpty<Event>(d_eventHeapSize)&&heapTop<Event>(d_eventHeap)->type == NewEllipse && heapTop<Event>(d_eventHeap)->distance - dOfFirstLevel < WXN_EPSILON ){
					cnt++;
					//if( cnt > 3 ) break;
					EllipseArc* arc;
					arc = &d_arcsPool[d_arcsPoolLength[0]];
					d_arcsPoolLength[0]++;

					arc->initial(d_center[0],heapTop(d_eventHeap)->site,NULL,NULL);
					if( headArc == NULL ){
						headArc = arc;
					}
					ArcHeapEle tmpArcHeapEle(arc->GetRotateAngle(),arc);
					heapInsert<ArcHeapEle>(tmpArcHeapEle,d_arcHeap,d_arcHeapSize);
					heapPop<Event>(lastEvent,d_eventHeap,d_eventHeapSize);
				}
				heapInsert<Event>(lastEvent,d_eventHeap,d_eventHeapSize);
			}else if( heapTop<Event>(d_eventHeap)->type == NewEllipse ){
				float angleSite = atan2f(heapTop<Event>(d_eventHeap)->site.y-d_center[0].y,heapTop<Event>(d_eventHeap)->site.x-d_center[0].x);
				EllipseArc *arc = headArc;
				bool find = false;
				bool fEntireCircle = (headArc->leftActiveEdge == NULL);
				bool fCloseToRight = false;
				bool fCloseToLeft = false;
				cuPrintf("-------------angleSite: %f\n",angleSite);
				//do{
				//	Pont angles = arc->geta
				//}

			}


		}
		Event tmp;
		//__shared__ Event d_eventHeap[EVENT_HEAP_SIZE];
		heapInitial<Event>(d_eventHeap,d_eventHeapSize);
		for(int i = 0; i < EVENT_HEAP_SIZE-10;++i){
			tmp.distance = i % 100;
			heapInsert<Event>(tmp,d_eventHeap,d_eventHeapSize);
		}
		for(int i = 0; i < EVENT_HEAP_SIZE-10;++i){
			heapPop<Event>(tmp,d_eventHeap,d_eventHeapSize);
			if( i % (EVENT_HEAP_SIZE / 10 ) == 0 ) cuPrintf( "%.2f " , tmp.distance );
		}
	}
	//cuPrintf("\n");
	//cuPrintf("time %f\n" , (float)(clock()-t1)/(float)1000000 );
	//cuPrintf("\n");

	//printf("Kernel : A[%d]=%f,B[%d]=%f,C[%d]=%f\n",i,A[i],i,B[i],i,C[i]);
}


__device__ void removeOutofdateEvents(Event* d_eventHeap,int& d_eventHeapSize)
{
	while( !heapEmpty<Event>(d_eventHeapSize) ){
		if( heapTop<Event>(d_eventHeap)->distance > FLT_MAX ){
			heapPop<Event>(d_eventHeap,d_eventHeapSize);
		}else if( heapTop<Event>(d_eventHeap)->type == NewEllipse ){
			break;
		}

		//Ellipse Vanish
		if( heapTop<Event>(d_eventHeap)->leftActiveEdge == heapTop<Event>(d_eventHeap)->arc->leftActiveEdge && heapTop<Event>(d_eventHeap)->rightActiveEdge == heapTop<Event>(d_eventHeap)->arc->rightActiveEdge ){
			break;
		}else{
			heapPop<Event>(d_eventHeap,d_eventHeapSize);
		}
	}

}
#endif